/*
 * Projekt: GMU 2013 - Canny hranovy detektor OpenCL/CUDA
 * Autori: Pavel Nevrala - xnevra00 (xnevra00@stud.fit.vutbr.cz)
 *         Martin Jurak - xjurak02 (xjurak02@stud.fit.vutbr.cz)
 */
#include "GaussianFilterCL.h"


GaussianFilterCL::GaussianFilterCL(cl_uint maxWorkGroupP, GaussianMask *mask, cl_context *contextP, cl_command_queue *commandQueueP, cl_program *programP): 
	CannyPartsBase(maxWorkGroupP), gaussianMask(mask), context(contextP), commandQueue(commandQueueP), program(programP)
{
}


GaussianFilterCL::~GaussianFilterCL(void)
{
}


/*
 * Provede nastaveni OpenCl a kernelu pro odstraneni sumu Gaussovym separabilnim filtrem.
 * @image - vstupni obraz
 * @width - sirka obrazu
 * @height - vyska obrazu	
 * @return - doba vypoctu na GPU
 */
double GaussianFilterCL::run(cl_float *image, cl_uint width, cl_uint height)
{
	cl_int clErr;
	cl_uint maskSize = gaussianMask->size();

	//vytvoreni bufferu
	cl_mem cl_in = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_in, CL_TRUE,	0, width * height * sizeof(cl_float), image, 0, NULL, NULL);
	cl_mem cl_mask = clCreateBuffer(*context, CL_MEM_READ_ONLY, maskSize * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_mask, CL_TRUE,	0, maskSize * sizeof(cl_float), gaussianMask->getGaussianMaskCL(), 0, NULL, NULL);
	
	cl_mem cl_outX = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_outY = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);

	//ziskani kernelu pro konvoluci ve smeru X
	kernelGaussX = clCreateKernel(*program, "convolutionX", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_result" );
	
   	clErr = clSetKernelArg(kernelGaussX, 0, sizeof(cl_mem), &cl_in);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelGaussX, 1, sizeof(cl_mem), &cl_outX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelGaussX, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelGaussX, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelGaussX, 4, sizeof(cl_float), &cl_mask);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelGaussX, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");	

	size_t GWS[2], LWS[2];
	GWS[0]=maxWorkGroup;
	GWS[1]=height;
	LWS[0]=maxWorkGroup;
	LWS[1]=1;

	cl_event event_gaussX;
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelGaussX, 2, NULL, GWS, LWS, 0, NULL, &event_gaussX);
	errorController.CheckOpenCLError(clErr, "enquekernel");

// ========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	//ziskani kernelu pro konvoluci ve smeru Y
	kernelGaussY = clCreateKernel(*program, "convolutionY", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_result" );

   	clErr = clSetKernelArg(kernelGaussY, 0, sizeof(cl_mem), &cl_outX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelGaussY, 1, sizeof(cl_mem), &cl_outY);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelGaussY, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelGaussY, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelGaussY, 4, sizeof(cl_float), &cl_mask);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelGaussY, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");		

	GWS[0]=width;
	GWS[1]=maxWorkGroup;
	LWS[0]=1;
	LWS[1]=maxWorkGroup;

	cl_event event_gaussY;
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelGaussY, 2, NULL, GWS, LWS, 1, &event_gaussX, &event_gaussY);
	errorController.CheckOpenCLError(clErr, "enquekernel");

	clErr = clWaitForEvents(1, &event_gaussY);
	errorController.CheckOpenCLError(clErr, "waitforevents Gauss");
	
	// zjisteni doby vypoctu
	clGetEventProfilingInfo(event_gaussX, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
	clGetEventProfilingInfo(event_gaussY, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
	total_time = (double)(time_end - time_start);	

	// zapis zpracovaneho obrazu
	outImage = new cl_float[width * height];
	memset(outImage, 0, width * height * sizeof(cl_float));
	clErr = clEnqueueReadBuffer (*commandQueue, cl_outY, CL_TRUE, 0, width * height * sizeof(cl_float), outImage, 0, NULL, NULL);
	errorController.CheckOpenCLError(clErr, "enqueread buffer Gauss");	

	// Uvolneni zdroju
    clErr = clReleaseKernel(kernelGaussX);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel gaussX.");
    clErr = clReleaseKernel(kernelGaussY);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel gaussY.");
	clErr = clReleaseMemObject(cl_in);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject input");    
    clErr = clReleaseMemObject(cl_outX);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject gaussX");
    clErr = clReleaseMemObject(cl_outY);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject gaussY");	
    clErr = clReleaseMemObject(cl_mask);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject mask");

	return total_time / 1000000.0;
}
